/*
 *  clhelper.c
 *  Sandbox
 *
 *  Created by Joshua Lewis on 10/28/10.
 *  Routines for reading system info and running OpenCL code.
 *
 */

#include "clhelper.h"


void getOpenCLInfo(){
	cl_int err;
	
	//get CPU info, up to 16 CPUs
	CPU_FOUND = clGetDeviceIDs(NULL,CL_DEVICE_TYPE_CPU,16,cpu,&n_CPUs);
	assert(CPU_FOUND == CL_SUCCESS);
	
	GPU_FOUND = clGetDeviceIDs(NULL,CL_DEVICE_TYPE_GPU,16,gpu,&n_GPUs);
	
	CPUs=malloc(sizeof(CLDevice)*n_CPUs);
	GPUs=malloc(sizeof(CLDevice)*n_GPUs);
	cl_char vendor[1024];
	cl_char devicename[1024];
	cl_uint available;
	size_t vendorsize,namesize,rsize;
	cl_ulong localsize,cachesize,cacheline,globalsize;
	
	for (int i=0; i<n_CPUs; i++) {
		err = clGetDeviceInfo(cpu[i],CL_DEVICE_NAME,sizeof(devicename),&devicename,&namesize);
		err = clGetDeviceInfo(cpu[i],CL_DEVICE_VENDOR,sizeof(vendor),&vendor,&vendorsize);
		err = clGetDeviceInfo(cpu[i],CL_DEVICE_AVAILABLE,sizeof(available),&available,&rsize);
		err = clGetDeviceInfo(cpu[i],CL_DEVICE_LOCAL_MEM_SIZE,sizeof(cl_ulong),&localsize,&rsize);
		err = clGetDeviceInfo(cpu[i],CL_DEVICE_GLOBAL_MEM_SIZE,sizeof(cl_ulong),&globalsize,&rsize);
		err = clGetDeviceInfo(cpu[i],CL_DEVICE_GLOBAL_MEM_CACHE_SIZE,sizeof(cl_ulong),&cachesize,&rsize);
		err = clGetDeviceInfo(cpu[i],CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE,sizeof(cl_ulong),&cacheline,&rsize);
		printf("\n\nCPU %i:\n",i);
		strncpy(CPUs[i].devicename,(char*)&devicename,namesize);
		strncpy(CPUs[i].vendor,(char*)&vendor,vendorsize);
		CPUs[i].available=(unsigned int)available;
		CPUs[i].local_mem_size=(unsigned long int)localsize;
		CPUs[i].global_mem_size=(unsigned long int)globalsize;
		CPUs[i].global_cache_size=(unsigned long int)cachesize;
		CPUs[i].global_cacheline_size=(unsigned long int)cacheline;
		
		printf("Name: %s\nVendor: %s\nAvailable: %i\nLocal Mem: %lu\nGlobal Mem: %lu\nGlobal Cache: %lu\nCacheline Size: %lu",
           CPUs[i].devicename,CPUs[i].vendor,CPUs[i].available,CPUs[i].local_mem_size,CPUs[i].global_mem_size,
           CPUs[i].global_cache_size,CPUs[i].global_cacheline_size);
	}
	for (int i=0; i<n_GPUs; i++) {
		err = clGetDeviceInfo(gpu[i],CL_DEVICE_NAME,sizeof(devicename),&devicename,&namesize);
		err = clGetDeviceInfo(gpu[i],CL_DEVICE_VENDOR,sizeof(vendor),&vendor,&vendorsize);
		err = clGetDeviceInfo(gpu[i],CL_DEVICE_AVAILABLE,sizeof(available),&available,&rsize);
		err = clGetDeviceInfo(gpu[i],CL_DEVICE_LOCAL_MEM_SIZE,sizeof(cl_ulong),&localsize,&rsize);
		err = clGetDeviceInfo(gpu[i],CL_DEVICE_GLOBAL_MEM_SIZE,sizeof(cl_ulong),&globalsize,&rsize);
		err = clGetDeviceInfo(gpu[i],CL_DEVICE_GLOBAL_MEM_CACHE_SIZE,sizeof(cl_ulong),&cachesize,&rsize);
		err = clGetDeviceInfo(gpu[i],CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE,sizeof(cl_ulong),&cacheline,&rsize);
		printf("\n\nGPU %i:\n",i);
		strncpy(GPUs[i].devicename,(char*)&devicename,namesize);
		strncpy(GPUs[i].vendor,(char*)&vendor,vendorsize);
		GPUs[i].available=(unsigned int)available;
		GPUs[i].local_mem_size=(unsigned long int)localsize;
		GPUs[i].global_mem_size=(unsigned long int)globalsize;
		GPUs[i].global_cache_size=(unsigned long int)cachesize;
		GPUs[i].global_cacheline_size=(unsigned long int)cacheline;
		
		printf("Name: %s\nVendor: %s\nAvailable: %i\nLocal Mem: %lu\nGlobal Mem: %lu\nGlobal Cache: %lu\nCacheline Size: %lu \n\n",
           GPUs[i].devicename,GPUs[i].vendor,GPUs[i].available,GPUs[i].local_mem_size,GPUs[i].global_mem_size,
           GPUs[i].global_cache_size,GPUs[i].global_cacheline_size);
	}
}


char * load_program_source(const char *filename)
{ 
	
	struct stat statbuf;
	FILE *fh; 
	char *source; 
	
	fh = fopen(filename, "r");
	if (fh == 0)
		return 0; 
	
	stat(filename, &statbuf);
	source = (char *) malloc(statbuf.st_size + 1);
	fread(source, statbuf.st_size, 1, fh);
	source[statbuf.st_size] = '\0'; 
	
	return source; 
} 


float * row_to_col(float * data, int N, int D)
{
	float *out = (float *)malloc(N * D * sizeof(float));
	for(int c=0; c<N; c++)
	{
		for (int r=0; r<D; r++) {
			out[c*D+r]=data[c+r*N];
		}
	}
	return out;
}


unsigned int * zeros (int N){
	unsigned int * out=(unsigned int*)malloc(N*sizeof(unsigned int));
	for (int i=0; i<N; i++) {
		out[i]=0;
	}
	return out;
}


double * dzeros (int N){
	double * out=(double*)malloc(N*sizeof(double));
	for (int i=0; i<N; i++) {
		out[i]=0;
	}
	return out;
}


void initQueue(){
	
	//getOpenCLInfo();
	GPU_FOUND = clGetDeviceIDs(NULL,CL_DEVICE_TYPE_GPU,1,&testdevice,NULL);
	context=clCreateContext(0, 1,&testdevice, NULL, NULL, &err);
	queue=clCreateCommandQueue(context,testdevice, 0, &err);
}


void createProgram(const char * source){
	program=clCreateProgramWithSource(context,1,(const char **)&source,NULL,&err);
	err=clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
  assert(err == CL_SUCCESS);
}


void createKernel(char * name){
	kernel = clCreateKernel(program, name, &err);
  assert(err == CL_SUCCESS);
}

double gpuPdist_CL(int N, int D, float *data){
	float * distance = (float *)malloc(N * N * sizeof(float));
	
	size_t local[2] = {16, 16};
	size_t global[2] = {32, 32};
	
	double elapsed;
	
	initQueue();
	const char * sourcefile = "gpuPdist.cl";
	char * programsource=load_program_source(sourcefile);
	createProgram(programsource);
	createKernel("gpuPdist");
	
	clinput = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(float) * N * D, NULL, NULL); //in: data matrix
	cloutput = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float) * N * N, NULL, NULL); //out: distance matrix (full)

	err = clEnqueueWriteBuffer(queue, clinput, CL_TRUE, 0, sizeof(float) * N * D, data, 0, NULL, NULL);
	assert(err == CL_SUCCESS);
	err = clFinish(queue);
	
	int i = 0;
	err = clSetKernelArg(kernel, i++, sizeof(cl_mem), &clinput); // in, data
	err = clSetKernelArg(kernel, i++, sizeof(cl_mem), &cloutput); // out, 
	err = clSetKernelArg(kernel, i++, sizeof(int), &N); // N
	err = clSetKernelArg(kernel, i++, sizeof(int), &D); // D
	err = clSetKernelArg(kernel, i++, 256 * sizeof(float), NULL); // Xfloat
	err = clSetKernelArg(kernel, i++, 256 * sizeof(float), NULL); // Yfloat

	assert(err == CL_SUCCESS);
	err=clFinish(queue);
	
	startTimer();
	err = clEnqueueNDRangeKernel(queue, kernel, 2, NULL, global, local, 0, NULL, NULL);
	assert(err == CL_SUCCESS);
	err = clFinish(queue);
	elapsed = stopTimer();
	
	err = clEnqueueReadBuffer(queue, cloutput, CL_TRUE, 0, sizeof(float)* N * N, distance, 0, NULL, NULL);
	assert(err == CL_SUCCESS);
	
//	for (int i = 0; i < N; i++) {
//		for (int j = 0; j < N; j++)
//			printf("%f \t", distance[i * N + j]);
//		printf("\n");
//	}

	
	free(distance);
	
	clReleaseMemObject(clinput);
	clReleaseMemObject(cloutput);
	clReleaseProgram(program);
	clReleaseKernel(kernel);
	clReleaseCommandQueue(queue);
	clReleaseContext(context);
	
	return elapsed;
}

void testCL_prune(int N, float *distance){
	
	dendrite * tree = (dendrite *)malloc((N * (N - 1) / 2) * sizeof(dendrite));
	for(int i = 0; i < N - 1; i++)
		for(int j = i + 1; j < N; j++) {
			tree[j * (j - 1) / 2 + i].distance = distance[j * (j - 1) / 2 + i];
			tree[j * (j - 1) / 2 + i].i = i;
			tree[j * (j - 1) / 2 + i].j = j;
		}
	
	psort(tree, N * (N - 1) / 2, sizeof(dendrite), compareDendrite);
	
	int * sorted = malloc(sizeof(int)* N * (N-1));
	for (int i = 0; i < (N*(N-1))/2; i++) {
		sorted[2*i] = tree[i].i;
		sorted[2*i + 1] = tree[i].j;
	}
	
//	for (int i = 0; i < (N*(N-1)/2); i++)
//		printf("%i %i \t", sorted[2*i], sorted[2*i+1]);	
//	printf("\n\n");
	
	int numWG = 4;
	size_t local = 16; //number of threads
	size_t global = numWG * (int)local;
  
  int baseSize = 128;
  int branchSize = N * (N - 1) / (2 * numWG);
	int twigSize =32;
  
  dendrite *base = malloc(sizeof(dendrite) * (N - 1));;
  int *baseStart = malloc(sizeof(int) * baseSize * 2);;
  dendrogram(N, distance, base);
  
  for(int i = 0; i < baseSize; i++) {
    baseStart[i * 2] = base[i].i;
    baseStart[i * 2 + 1] = base[i].j;
    printf("%i %i\t\t", baseStart[i * 2], baseStart[i * 2 + 1]);
  }
 printf("\n\n\n");
    
  int * sorted_result = malloc(sizeof(int) * numWG * (N - baseSize - 1) * 2);
  
	initQueue();
	const char * sourcefile = "dendrogram.cl";
	char * programsource=load_program_source(sourcefile);
	createProgram(programsource);
	createKernel("dendrogram");
	
	clinput = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(int) * N * (N - 1), NULL, NULL); //branch gets written
	cloutput = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(int) * numWG * (N - baseSize - 1) * 2, NULL, NULL);
  clbasestart = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(int) * baseSize * 2, NULL, NULL);
	err = clEnqueueWriteBuffer(queue, clinput, CL_TRUE, 0, sizeof(int) * N * (N - 1), sorted, 0, NULL, NULL);
  err = clEnqueueWriteBuffer(queue, clbasestart, CL_TRUE, 0, sizeof(int) * baseSize * 2, baseStart, 0, NULL, NULL);
	assert(err == CL_SUCCESS);
	err = clFinish(queue);
	
  int i = 0;
	err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &clinput); // sorted
	err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &cloutput); // result
	err = clSetKernelArg(kernel, 2, sizeof(cl_mem), &clbasestart); // baseStart
	err = clSetKernelArg(kernel, 3, sizeof(int), &baseSize); // baseStartSize
	err = clSetKernelArg(kernel, 4, sizeof(int), &N); // N
	err = clSetKernelArg(kernel, 5, sizeof(int), &branchSize); // banchSize
	err = clSetKernelArg(kernel, 6, sizeof(int), &twigSize); // twigSize
    err = clSetKernelArg(kernel, 7, sizeof(int) * (N - 1) * 2, NULL); // base
    err = clSetKernelArg(kernel, 8, sizeof(int) * (N - baseSize - 1) * 2, NULL); // baseClean
	err = clSetKernelArg(kernel, 9, (int)local * twigSize * sizeof(int) * 2, NULL); // branch
	err = clSetKernelArg(kernel, 10, (int)local * twigSize * sizeof(int) * 2, NULL); // branchClean
	err = clSetKernelArg(kernel, 11, (int)local * sizeof(int), NULL);
	assert(err == CL_SUCCESS);
	err=clFinish(queue);
	err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, &local, 0, NULL, NULL);
	assert(err == CL_SUCCESS);
	err = clFinish(queue);
	
	err = clEnqueueReadBuffer(queue, cloutput, CL_TRUE, 0, sizeof(int)* numWG * (N - baseSize - 1) * 2, sorted_result, 0, NULL, NULL);
	assert(err == CL_SUCCESS);
	
	for (int i = 0; i < numWG * (N - baseSize - 1); i++) {
    if(i % (N - baseSize - 1) == 0)
      printf("\n\n\n");
		printf("%i %i\t\t", sorted_result[2*i], sorted_result[2*i +1]);
 }
	
  free(tree);
  free(base);
  free(baseStart);
  free(sorted);
  free(sorted_result);
  
	clReleaseMemObject(clinput);
	clReleaseMemObject(cloutput);
	clReleaseProgram(program);
	clReleaseKernel(kernel);
	clReleaseCommandQueue(queue);
	clReleaseContext(context);
}

double testCL(unsigned int N, unsigned int D, unsigned int blocksize, int n_tests, float * data, float *comparison){
	float * result;
  double t_out, t_comp;
	unsigned int *temp;
	temp=zeros(N/blocksize);
	size_t global, local;
	cl_mem cross_counter, row_counter;
	result=malloc(sizeof(float)*N*(N-1)/2);
	initQueue();
	const char * sourcefile="distanceOne.cl";
	char * programsource=load_program_source(sourcefile);
	createProgram(programsource);
	createKernel("distanceOne");
    unsigned int rowstart=0;
	
	clinput = clCreateBuffer(context,  CL_MEM_READ_ONLY, sizeof(float)*N*D, NULL, NULL);
	cloutput = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float)*N*(N-1)/2, NULL, NULL);	
    cross_counter = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(unsigned int)*N/blocksize,NULL,NULL);
	row_counter = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(unsigned int),NULL,NULL);

	
	err = clEnqueueWriteBuffer(queue, clinput, CL_TRUE, 0, sizeof(float)*N*D, data, 0, NULL, NULL);
	err = clEnqueueWriteBuffer(queue, cross_counter, CL_TRUE, 0, sizeof(unsigned int)*N/blocksize, temp, 0, NULL, NULL);
	err = clEnqueueWriteBuffer(queue, row_counter, CL_TRUE, 0, sizeof(unsigned int), &rowstart, 0, NULL, NULL);
	
	
	
	err = clFinish(queue);
	
	err=clGetKernelWorkGroupInfo(kernel, testdevice, CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), &localsize, NULL);	
	
	err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &clinput);
	err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &cloutput);
	err = clSetKernelArg(kernel, 2, sizeof(float)*D*blocksize, NULL);
	err = clSetKernelArg(kernel, 3, sizeof(float)*D*blocksize, NULL);
  err = clSetKernelArg(kernel, 4, sizeof(float)*blocksize*blocksize, NULL);
	err = clSetKernelArg(kernel, 5, sizeof(unsigned int), &N);
	err = clSetKernelArg(kernel, 6, sizeof(unsigned int), &D);
   
  
  
//  err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &clinput);
//	err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &cloutput);
//	err = clSetKernelArg(kernel, 2, sizeof(cl_mem), &row_counter);
//	err = clSetKernelArg(kernel, 3, sizeof(cl_mem), &cross_counter);
//	err = clSetKernelArg(kernel, 4, sizeof(float)*D*blocksize, NULL);
//	err = clSetKernelArg(kernel, 5, sizeof(float)*D*blocksize, NULL);
//	err = clSetKernelArg(kernel, 6, sizeof(unsigned int), NULL);
//	err = clSetKernelArg(kernel, 7, sizeof(unsigned int), &N);
//	err = clSetKernelArg(kernel, 8, sizeof(unsigned int), &D);
//	err = clSetKernelArg(kernel, 9, sizeof(unsigned int), &blocksize);
  
	
	err = clFinish(queue);
	
  
	global=(size_t)N*(N/blocksize+1)/2;
	local=(size_t)blocksize;
  
  //int nwg = global / local;
	
	
	startTimer();
	err = clEnqueueNDRangeKernel(queue, kernel,1, NULL, &global,&local, 0, NULL, NULL);
	err = clFinish(queue);
	t_out = stopTimer();
	
	err = clEnqueueReadBuffer(queue, cloutput, CL_TRUE, 0, sizeof(float)*N*(N-1)/2, result, 0, NULL, NULL);
    //err = clEnqueueReadBuffer(queue, cross_counter, CL_TRUE, 0, sizeof(float)*N/blocksize, temp, 0, NULL, NULL);
	
	
//	int disterror=0;
//	for (int i=0; i<N*(N-1)/2; i+=1) {
//		
//		if(fabs(result[i]-comparison[i])>.001)
//		{
//			//printf("\n%f === %f, %d\n",result[i],comparison[i], i);
//			disterror++;
//		}
//	}
	
  /*
  for(int i = 0; i < N; i++) {
    for(int j = 0; j < N; j++) {
      if(i == j)
        printf("0\t");
      else {
        if(fabs(result[utndidx(fmin(i, j), fmax(i, j))]-comparison[utndidx(fmin(i, j), fmax(i, j))])>.001)
          printf("1\t");
        else
          printf("0\t");
      }
    }
    printf("\n");
  }
  */
  
  //for(int i = 0; i < N / blocksize; i++)
  //  printf("%d\t",temp[i]);
  //printf("\n\n");
  
//	if (disterror>0) {
//		printf("\n\n%d errors with:\nN = %i\nD = %i\nblocksize = %i\n\n",disterror,N,D,blocksize);
//	}
//   
//	t_out=999999;
//	t_comp=0;
//	if (disterror==0) {
//		
//	
//	
//	for (int i=0; i<n_tests; i++) {
		
		//startTimer();
    
		//err = clEnqueueWriteBuffer(queue, clinput, CL_TRUE, 0, sizeof(float)*N*D, data, 0, NULL, NULL);
//		err = clEnqueueWriteBuffer(queue, cross_counter, CL_TRUE, 0, sizeof(unsigned int)*N/blocksize, temp, 0, NULL, NULL);
//		err = clEnqueueWriteBuffer(queue, row_counter, CL_TRUE, 0, sizeof(unsigned int), &rowstart, 0, NULL, NULL);
//		
//		
//		err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &clinput);
//		err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &cloutput);
//		err = clSetKernelArg(kernel, 2, sizeof(cl_mem), &row_counter);
//		err = clSetKernelArg(kernel, 3, sizeof(cl_mem), &cross_counter);
//		err = clSetKernelArg(kernel, 4, sizeof(float)*D*blocksize, NULL);
//		err = clSetKernelArg(kernel, 5, sizeof(float)*D*blocksize, NULL);
//		err = clSetKernelArg(kernel, 6, sizeof(unsigned int), NULL);
//		err = clSetKernelArg(kernel, 7, sizeof(unsigned int), &N);
//		err = clSetKernelArg(kernel, 8, sizeof(unsigned int), &D);
//		err = clSetKernelArg(kernel, 9, sizeof(unsigned int), &blocksize);
		
		
//		err = clEnqueueNDRangeKernel(queue, kernel,1, NULL, &global,&local, 0, NULL, NULL);
//		err = clFinish(queue);
//		if (err!=0) {
//			printf("PROBLEM\n\n");
//		}
//		
//		
//		t_comp=stopTimer();
//		if(t_comp<t_out){
//			t_out=t_comp;
//		}
//		
//	}
	
//	}
	clReleaseMemObject(clinput);
	clReleaseMemObject(cloutput);
	clReleaseMemObject(cross_counter);
	clReleaseMemObject(row_counter);
	clReleaseProgram(program);
	clReleaseKernel(kernel);
	clReleaseCommandQueue(queue);
	clReleaseContext(context);
	return t_out;
}


void runTests(int N, int D, int n_conds,float * data)
{
	float * data_cm;
	float *comparison=malloc(sizeof(float)*N*(N-1)/2);
	data_cm=row_to_col(data, N, D);
  startTimer();
	distance3(N, D, data_cm, comparison);
  double cputime = stopTimer();
  printf("CPU Time: %f\n\n",cputime);
  
  
	double * results;
	unsigned int * bsizes;
	results=dzeros(n_conds);
	bsizes=zeros(n_conds);
	
	for (int i=0; i<n_conds;i++) {
		bsizes[i]=8 * pow(2, i);
		results[i]=testCL(N,D,bsizes[i], 5,data, comparison);
	}
	
	for (int i=0; i<n_conds; i++) {
		printf("Blocksize: %i\nTime: %f\n\n",bsizes[i],results[i]);
	}
  
  free(comparison);
}
